#include "compress_gpu.h"

__device__ void d_binary_encoding(uchar4* data, uchar4* out)
{
	uchar4 d, o;
	uint32 i;
	o.x = 0;
	o.y = 0;
	o.z = 0;
	o.w = 0;	

	d = data[0];
	o.x <<= 2;
	o.x |= d.x;
	o.x <<= 2;
	o.x |= d.y;
	o.x <<= 2;
	o.x |= d.z;
	o.x <<= 2;
	o.x |= d.w;

	d = data[1];
	o.y <<= 2;
	o.y |= d.x;
	o.y <<= 2;
	o.y |= d.y;
	o.y <<= 2;
	o.y |= d.z;
	o.y <<= 2;
	o.y |= d.w;

	d = data[2];
	o.z <<= 2;
	o.z |= d.x;
	o.z <<= 2;
	o.z |= d.y;
	o.z <<= 2;
	o.z |= d.z;
	o.z <<= 2;
	o.z |= d.w;

	d = data[3];
	o.w <<= 2;
	o.w |= d.x;
	o.w <<= 2;
	o.w |= d.y;
	o.w <<= 2;
	o.w |= d.z;
	o.w <<= 2;
	o.w |= d.w;

	out[0] = o;
}

__global__ void binary_encoding_kernel(uint8* data, uint32 dataSize, uint8* out)
{
	uint32 i = blockIdx.x * blockDim.x + threadIdx.x;
	uint32 dataSize4 = dataSize >> 2;
	uchar4 *data4 = (uchar4 *)data;
	uchar4 *out4 = (uchar4 *)out;
	uint32 outPosIndex = 0;
	uint32 curPos;
	while(i*4 < dataSize4) {
		d_binary_encoding(data4 + i*4, out4 + i);		
		i += blockDim.x * gridDim.x;
	}
}

void gpu_binary_encoding(uint8* d_data, uint32 dataSize, uint8* d_out)
{
	INIT_TIMER;
	START_TIMER;

	uint32 thread_num = 128, block_num = 512;
	
	binary_encoding_kernel<<<block_num, thread_num>>>(d_data,dataSize, d_out);
	cutilCheckMsg("binary_encoding_kernel");
	cutilSafeCall(cudaThreadSynchronize());

	END_TIMER;
	PRINT_TIMER_SEC("binary_encoding_kernel ");
}

void gpu_binary_encoding_test(uint8* data, uint32 dataSize, uint8* out)
{
	uint8* d_data; 
	uint8* d_out;
	uint32 dataSize1 = dataSize; 
	if(dataSize % 16 > 0) 
		dataSize1 = (dataSize / 16 + 1 ) * 16;

	INIT_TIMER;
	START_TIMER;
	cutilSafeCall(cudaMalloc(&d_data,  dataSize1));
	cutilSafeCall(cudaMemset(d_data, 0, dataSize1));
	cutilSafeCall(cudaMemcpy(d_data, data, dataSize, cudaMemcpyHostToDevice));
	cutilSafeCall(cudaMalloc(&d_out, dataSize1 / 4));
	cutilSafeCall(cudaMemset(d_out, 0, dataSize1 / 4));
	END_TIMER;
	PRINT_TIMER_SEC("malloc time ");
	
	START_TIMER;

	uint32 thread_num = 128, block_num = 512;
	
	binary_encoding_kernel<<<block_num, thread_num>>>(d_data, dataSize1, d_out);
	cutilCheckMsg("binary_encoding_kernel");
	cutilSafeCall(cudaThreadSynchronize());

	END_TIMER;
	PRINT_TIMER_SEC("binary_encoding_kernel ");

	cutilSafeCall(cudaMemcpy(out, d_out, ceil((double)dataSize/4), cudaMemcpyDeviceToHost));

	cutilSafeCall(cudaFree(d_data));
	cutilSafeCall(cudaFree(d_out));
}

